// simple.cu

// Copyright (C) 2010  traxex

// Author: 	traxex <traxex@yahoo.cn>
// Created: 	10 Aug 2010
// Version: 	1.0

#include "cuda/unittest.h"

//Oh, SHIT! What a bad code I make.

#ifndef OPNAME
const int 		N = 4096*4096;
typedef vector<float>	Vector;

#	define 	OPNAME	add
#	define 	OP 	+
#	include	"simple.cu"

#	define 	OPNAME	sub
#	define 	OP 	-
#	include	"simple.cu"

#	define 	OPNAME	mul
#	define 	OP 	*
#	include	"simple.cu"

#	define 	OPNAME	div
#	define 	OP 	/
#	include	"simple.cu"

void performance(){
    performance_add();
    performance_sub();
    performance_mul();
    performance_div();
}

bool unittest(){

    float 	ha[N], hb[N], hc[N];
    Vector	da(N), db(N), dc(N);

    for( int i = 0; i < N; ++i){
	ha[i] = i*i;
	hb[i] = i;
    }

    da << ha;
    db << hb;

    for( int i = 0; i < N; ++i){	hc[i] = ha[i]+hb[i];    }
    simple_kernel_add<<<cusize(N)>>>( dc, da, db);

    float		dhc[N];
    dc >> dhc;
    
    _vview( "%.2f, ", hc, 10);
    _vview( "%.2f, ", dhc, 10);
    return vcmp( hc, dhc, N);
}
#else


__global__ void CAT(simple_kernel_,OPNAME)( Vector dc, Vector da, Vector db){
    if( _vi >= N) return;
    dc(_vi) = da(_vi) OP db(_vi);
}

void CAT(performance_,OPNAME)( ){
    double t1 = timenow();
    Vector	da(N), db(N), dc(N);    
    LOOP CAT(simple_kernel_,OPNAME)<<<cusize(N)>>>( dc, da, db);
    cuda::sync();
    double t2 = timenow();
    printf( "c=a%sb	%f sec, %d float, %f Gflops\n", STR(OP), (t2-t1)/LOOPN, N, N/1e9*LOOPN/(t2-t1));
}

#undef OP
#undef OPNAME
#endif
